Skip to content

Implement the new tuning API for DeviceRadixSort #6767

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom
Open
bernhardmgruber wants to merge 25 commits into
base: main
Choose a base branch
from
Open

Implement the new tuning API for DeviceRadixSort #6767

bernhardmgruber wants to merge 25 commits into from

Conversation

@bernhardmgruber
Copy link
Contributor

@bernhardmgruber bernhardmgruber commented Nov 24, 2025 •
edited
Loading

PR is not fully ready yet, but ready to take a first round of review.

  • cub.test.device.radix_sort_keys.lid_0.key_bits_16 passes
  • CCCL.C tests pass
  • Retain the if constexpr on the onesweep algorithm in the dispatcher
  • implement policy_selector_from_hub
  • Drop _policy
  • Merge before: Rename new tuning API policies and fix MSVC warning #7103
  • tests for supplying a custom policy hub
  • SASS tests

Fixes: #6676

@github-project-automation github-project-automation bot moved this to Todo in CCCL Nov 24, 2025
@bernhardmgruber bernhardmgruber changed the title Implement the new tuning API for DeviceRadixSort Implement the new tuning API for DeviceRadixSort Nov 24, 2025
@cccl-authenticator-app cccl-authenticator-app bot moved this from Todo to In Progress in CCCL Nov 24, 2025
miscco
miscco approved these changes Nov 25, 2025
View reviewed changes
Copy link
Contributor

@miscco miscco left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I love how much this cleans everything up

cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh
Comment on lines +134 to +176
_CCCL_API constexpr friend bool
operator!=(const radix_sort_histogram_policy& lhs, const radix_sort_histogram_policy& rhs)
{
return !(lhs == rhs);
}
Copy link
Contributor

@miscco miscco Nov 25, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Nitpick: We could guard all operator!= on #if _CCCL_STD_VER <= 2017

Copy link
Contributor Author

@bernhardmgruber bernhardmgruber Nov 25, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Or we just wait until C++20 and refactor.

@bernhardmgruber bernhardmgruber mentioned this pull request Nov 25, 2025
Open
10 tasks
@bernhardmgruber bernhardmgruber force-pushed the tuning_radix_sort branch 2 times, most recently from 0d56d8d to 56c437d Compare November 25, 2025 12:09
@bernhardmgruber
Copy link
Contributor Author

bernhardmgruber commented Nov 25, 2025

/ok to test 56c437d

@github-actions

This comment has been minimized.

Sign in to view
@bernhardmgruber bernhardmgruber force-pushed the tuning_radix_sort branch 3 times, most recently from 8ad87dc to 800b936 Compare January 11, 2026 19:03
@bernhardmgruber
Copy link
Contributor Author

bernhardmgruber commented Jan 11, 2026

/ok to test 800b936

@github-actions

This comment has been minimized.

Sign in to view
@NVIDIA NVIDIA deleted a comment from copy-pr-bot bot Jan 11, 2026
@bernhardmgruber bernhardmgruber mentioned this pull request Jan 12, 2026
Open
30 tasks
@copy-pr-bot
Copy link
Contributor

copy-pr-bot bot commented Jan 12, 2026

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

@bernhardmgruber
Copy link
Contributor Author

bernhardmgruber commented Jan 12, 2026

/ok to test 7e722a2

@github-actions

This comment has been minimized.

Sign in to view
@bernhardmgruber
Copy link
Contributor Author

bernhardmgruber commented Jan 12, 2026

/ok to test 6cefe64

@bernhardmgruber
Copy link
Contributor Author

bernhardmgruber commented Jan 12, 2026

/ok to test 74089a2

@github-actions

This comment has been minimized.

Sign in to view
@cccl-authenticator-app cccl-authenticator-app bot moved this from In Progress to In Review in CCCL Jan 13, 2026
@github-actions
Copy link
Contributor

github-actions bot commented Jan 13, 2026

😬 CI Workflow Results

🟥 Finished in 3h 19m: Pass: 72%/93 | Total: 2d 23h | Max: 2h 38m | Hits: 62%/74480

See results here.

@bernhardmgruber
Copy link
Contributor Author

bernhardmgruber commented Jan 13, 2026 

/home/coder/cccl/lib/cmake/cub/../../../cub/cub/detail/arch_dispatch.cuh:65:34: internal compiler error: side-effects element in no-side-effects CONSTRUCTOR
   static constexpr policy_t all_policies[sizeof...(Is)]      = {PolicySelector{}(all_arches[Is])...};
                                  ^~~~~~~~~~~~

Nice! I have never seen this one!

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@miscco miscco miscco approved these changes

@shwina shwina Awaiting requested review from shwina shwina is a code owner automatically assigned from NVIDIA/cccl-benchmark-codeowners

@elstehle elstehle Awaiting requested review from elstehle elstehle is a code owner automatically assigned from NVIDIA/cccl-cub-codeowners

At least 1 approving review is required to merge this pull request.

Assignees

No one assigned

Labels

None yet

Projects

CCCL
Status: In Review

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

Implement the new tuning API for cub::DeviceRadixSort

2 participants

@bernhardmgruber @miscco